CUDA: enable the CUB path on HIP via hipCUB by Geramy · Pull Request #26592 · ggml-org/llama.cpp

Geramy · 2026-08-04T17:12:57Z

Overview

Since the CUB-based ARGSORT/TOP_K paths were never running on a HIP based device, I have enabled them and added a few missing hip graph functions. Originally those ops fell back to the shared-memory bitonic sort. That sort uses one thread per element, so it is limited to rows of 1024 elements; anything wider was reported as unsupported and ran on the CPU.

Additional information

I have added stream-capture functions missing from vendors/hip.h which is required in CUB argsort.

DeepSeek-V4-Flash UD-IQ1_S — 284B MoE, 78 GB

test	depth	main (t/s)	PR (t/s)	Δ (t/s)	change
`pp512` _{prompt processing}	0	162.28 _{± 4.83}	163.17 _{± 3.21}	+0.89	+0.55% ± 3.59
	2048	156.02 _{± 1.28}	156.67 _{± 0.85}	+0.65	+0.41% ± 0.99
	4096	139.08 _{± 0.48}	146.72 _{± 1.66}	+7.64	+5.49% ± 1.25
	6192	132.30 _{± 0.88}	140.59 _{± 0.99}	+8.29	+6.26% ± 1.03
`tg32` _{token generation}	0	16.05 _{± 0.15}	16.11 _{± 0.11}	+0.06	+0.40% ± 1.16
	2048	14.24 _{± 0.13}	14.34 _{± 0.01}	+0.11	+0.74% ± 0.95
	4096	14.64 _{± 0.25}	15.55 _{± 0.12}	+0.91	+6.22% ± 2.00
	6192	14.68 _{± 0.08}	15.65 _{± 0.15}	+0.96	+6.56% ± 1.20

Qwen3.6-35B-A3B UD-Q4_K_XL

test	depth	main (t/s)	PR (t/s)	Δ (t/s)	change
`pp512` _{prompt processing}	0	2064.78 _{± 967.80}	2061.34 _{± 948.84}	-3.44	-0.17% ± 65.58
	2048	1890.68 _{± 868.16}	1884.03 _{± 856.98}	-6.65	-0.35% ± 64.41
	4096	1766.93 _{± 781.34}	1748.30 _{± 774.43}	-18.63	-1.05% ± 61.93
	6192	1644.69 _{± 716.90}	1636.84 _{± 703.89}	-7.84	-0.48% ± 60.94
`tg32` _{token generation}	0	54.66 _{± 4.56}	54.57 _{± 3.53}	-0.09	-0.17% ± 10.54
	2048	55.32 _{± 4.69}	55.28 _{± 4.62}	-0.04	-0.08% ± 11.90
	4096	55.09 _{± 4.89}	55.03 _{± 4.82}	-0.06	-0.11% ± 12.46
	6192	54.69 _{± 4.94}	54.61 _{± 4.91}	-0.08	-0.15% ± 12.72

Requirements

AMD GPU and 2* AMD GPUs to test tensor splitting.
110GB+ VRAM
Deepseek v4 flash

I have read and agree with the contributing guidelines Yes
AI usage disclosure: Yes - for research purposes.

@IMbackK @am17an @pwilkin

The CUB-based ARGSORT/TOP_K (and sum/mean/cumsum) paths were compiled out for HIP, so those ops fell back to the shared-memory bitonic sort. That sort uses one thread per element, so it is limited to rows of 1024 elements; anything wider was reported as unsupported and ran on the CPU. For deepseek4 that meant every lightning-indexer TOP_K above 1024 context went to the host, adding a graph split per layer. hipCUB provides the device-wide sorts these paths need - the rocPRIM backend has DeviceRadixSort, DeviceSegmentedRadixSort, DeviceSegmentedSort, DeviceReduce and DeviceScan - so include it through a small cub-compat shim that aliases the cub namespace, and define GGML_CUDA_USE_CUB for HIP as well. hipCUB has no DeviceTopK and does not define CCCL_*, so CUB_TOP_K_AVAILABLE and STRIDED_ITERATOR_AVAILABLE stay disabled there: TOP_K keeps using the argsort + copy path and argsort keeps using the init_offsets kernel instead of a strided iterator. vendors/hip.h gains the stream-capture defines that the CUB argsort path needs now that it is compiled for HIP.

fairydreaming · 2026-08-04T18:30:50Z

@ggml-org/ci Looks like CI hip workflows need to be updated to install hipcub dev package for this. Not sure about HIP Windows release, but I guess that's something to worry about later.

Neresco · 2026-08-04T18:33:38Z

So test this PR here with 4x gfx1200 9060xt and 1x gfx1151.
Slight increase in VRAM usage over the Prompt Processing (PP).
~400MB over 65k Context but i have not closely observed.
Two times my Window Manager crashed before i recognize the reason.
Oddly not the GPU's where the Monitors are plugged in had too much VRAM.

No crash at 68727 PP context.
It let it run further now.

Geramy · 2026-08-04T18:41:10Z

@fairydreaming isn't hipcub included in rocm?

fairydreaming · 2026-08-04T19:07:24Z

@fairydreaming isn't hipcub included in rocm?

@Geramy No idea, my adventure with AMD GPUs ended during Radeon R9 290 times. By the way, could you run some test-backend-ops test_top_k test cases with large shapes on AMD GPU with this PR to make sure there are no issues for huge tensors? For example these two:

test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {262144, 8192, 1, 1}, 1024));
test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {1048576, 512, 1, 1}, 2048));

Thanks!

Geramy · 2026-08-04T19:18:34Z

@fairydreaming from AMD GitHub page directly hipCUB is included with the ROCm Core SDK on Linux and Windows.
Your tests also passed, I'm moving this from draft to review.

TOP_K(type=f32,ne=[262144,8192,1,1],k=1024,ties=0): OK
TOP_K(type=f32,ne=[1048576,512,1,1],k=2048,ties=0): OK
11/11 tests passed
3/3 backends passed
OK

Neresco · 2026-08-04T19:27:06Z

Cline is working now including reasoning and cache with up to 172k tokens for me over rpc.

Edit: by 193k tokens ca. 600mb VRAM increase per GPU i can observe.
So avoid too tight fits for AMD GPU's.

fairydreaming · 2026-08-04T19:40:27Z

@fairydreaming from AMD GitHub page directly `hipCUB is included with the ROCm Core SDK on Linux and Windows.

Well for some reason hipCUB header files were missing in the CI container during compilation, I see that it installs Ubuntu rocblas-dev and hipblas-dev packages, so maybe they also need to install libhipcub-dev or something.

Geramy · 2026-08-04T19:46:21Z

@fairydreaming from AMD GitHub page directly `hipCUB is included with the ROCm Core SDK on Linux and Windows.

Well for some reason hipCUB header files were missing in the CI container during compilation, I see that it installs Ubuntu rocblas-dev and hipblas-dev packages, so maybe they also need to install libhipcub-dev or something.

Yeah perhaps it depends on the ROCm version we are on too. Who will or can take care of that?

fairydreaming · 2026-08-04T19:50:53Z

@fairydreaming from AMD GitHub page directly `hipCUB is included with the ROCm Core SDK on Linux and Windows.

Well for some reason hipCUB header files were missing in the CI container during compilation, I see that it installs Ubuntu rocblas-dev and hipblas-dev packages, so maybe they also need to install libhipcub-dev or something.

Yeah perhaps it depends on the ROCm version we are on too. Who will or can take care of that?

@Geramy I don't know, if no one fixes it I'll see what I can do tomorrow (going to sleep soon).

remeh · 2026-08-04T20:01:07Z

Confirmed to work on a Strix Halo: with this patch the tg speed stays correct even after >4-5k context (instead of having a big drop to 5 tok/s). 🎉

However, I'm wondering if the prefill speed doesn't get worse faster than with #26493, but that can be totally wrong, don't read too much into this.

fairydreaming · 2026-08-05T09:03:44Z

HIP quality check CI found some missing CUB function call return value error checks, so I added them. Compilation worked fine with added hipcub-dev package.

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>

Geramy · 2026-08-05T13:18:19Z

@fairydreaming thanks! I didn’t see any of the CIs run or get approved to run did you do a few manually in actions?

fairydreaming · 2026-08-05T13:42:12Z

@Geramy Yeah, in this PR I have to approve workflows manually to run and they reset after each new commit. Probably you can see them in Actions history (ran about 4 hours ago).

Now let's wait for someone from ggml-cuda to review and approve this.

kyuz0 · 2026-08-05T14:04:50Z

With that version of the file it compiles and works correctly, slightly faster than the previous PR (on a Strix Halo). There may be a better way to resolve this naming conflict, but the above works on my system.

Are you able to publish the numbers you get? On ROCm-7.14 I am not getting better performance.

matt23654 · 2026-08-05T14:11:01Z

With that version of the file it compiles and works correctly, slightly faster than the previous PR (on a Strix Halo). There may be a better way to resolve this naming conflict, but the above works on my system.

Are you able to publish the numbers you get? On ROCm-7.14 I am not getting better performance.

14.2 tok/s generation (this PR) vs 14 tok/s generation (bitonic PR) - this is not a rigorous benchmark like yours though, this is just a casual observation for a test prompt. Even a rigorous benchmark would not be directly comparable because I am running over RPC and CUDA (Host) + RPC (ROCm) crashes on master without a PR to provide ROCm-based topk.

Edit: Did you also try with TheRock-Nightly?

Geramy · 2026-08-05T14:13:52Z

With that version of the file it compiles and works correctly, slightly faster than the previous PR (on a Strix Halo). There may be a better way to resolve this naming conflict, but the above works on my system.

Are you able to publish the numbers you get? On ROCm-7.14 I am not getting better performance.

The original issue is DeepSeek 4 Flash falls back to cpu top-k causing huge performance drop. You have to benchmark above 4096 to see the drop on main but this drop was over multiple GPUs.

kyuz0 · 2026-08-05T14:14:48Z

The original issue is DeepSeek 4 Flash falls back to cpu top-k causing huge performance drop. You have to benchmark above 4096 to see the drop on main but this drop was over multiple GPUs.

Oh, you mean generating more than 4K, not depth?

@fairydreaming

…pecific files for cub, dropped redundent includes and added a hip version check for hip cub inclusion, moved around a few includes, and added a test @fairydreaming originally had me run.

Geramy · 2026-08-05T16:01:33Z

The original issue is DeepSeek 4 Flash falls back to cpu top-k causing huge performance drop. You have to benchmark above 4096 to see the drop on main but this drop was over multiple GPUs.

Oh, you mean generating more than 4K, not depth?

Right the original issue is when I believe token input goes over like 50xx it causes top-k to fall back to the CPU top-k / argsort function. There is a optimization to be had later one perhaps, I do notice we "define" top-k and argsort as the same backend function when top-k doesn't have to be sorted, I would think this would save some compute and a little bit of memory bandwidth but that's a discussion for another PR, and this is all cuda / hip related backend, not vulkan.

Geramy · 2026-08-05T16:20:56Z

With that version of the file it compiles and works correctly, slightly faster than the previous PR (on a Strix Halo). There may be a better way to resolve this naming conflict, but the above works on my system.

Are you able to publish the numbers you get? On ROCm-7.14 I am not getting better performance.

what do you get on 7.13 or 7.2.3 or are we at .4 now? just wondering if its a 7.14 issue or a hipcub issue, I can look into hipcub a little more and see if there is something I can do about the performance.

fairydreaming · 2026-08-05T17:14:47Z

Regarding the questionable performance improvements observed by some people I suppose it depends on your specific CPU/GPU combination and interface between them. For example on my machine (Epyc 9374F + RTX PRO 6000 Max-Q) I have:

CPU

$ ./bin/test-backend-ops -b CPU perf -o "TOP_K"
ggml_cuda_init: found 1 CUDA devices (Total VRAM: 97247 MiB):
  Device 0: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition, compute capability 12.0, VMM: yes, VRAM: 97247 MiB
Testing 2 devices

Backend 1/2: CUDA0
  Skipping
Backend 2/2: CPU
  Device description: AMD EPYC 9374F 32-Core Processor
  Device memory: 1160411 MB (1160411 MB free)

  TOP_K(type=f32,ne=[4096,1,1,1],k=1024,ties=0):               32768 runs -    38.01 us/run -       20 kB/run -    0.50 GB/s
  TOP_K(type=f32,ne=[8192,1,1,1],k=1024,ties=0):               24576 runs -    47.66 us/run -       36 kB/run -    0.72 GB/s
  TOP_K(type=f32,ne=[16384,1,1,1],k=1024,ties=0):              16384 runs -    61.80 us/run -       68 kB/run -    1.05 GB/s
  TOP_K(type=f32,ne=[32768,1,1,1],k=1024,ties=0):              16384 runs -    78.22 us/run -      132 kB/run -    1.61 GB/s
  TOP_K(type=f32,ne=[65000,1,1,1],k=1024,ties=0):              16384 runs -   108.26 us/run -      257 kB/run -    2.27 GB/s
  TOP_K(type=f32,ne=[200000,1,1,1],k=1024,ties=0):              8192 runs -   536.51 us/run -      785 kB/run -    1.40 GB/s
  TOP_K(type=f32,ne=[4096,512,1,1],k=1024,ties=0):               820 runs -  1464.57 us/run -    10240 kB/run -    6.67 GB/s
  TOP_K(type=f32,ne=[8192,512,1,1],k=1024,ties=0):               912 runs -  1985.42 us/run -    18432 kB/run -    8.85 GB/s
  TOP_K(type=f32,ne=[16384,512,1,1],k=1024,ties=0):                      482 runs -  2522.04 us/run -    34816 kB/run -   13.17 GB/s
  TOP_K(type=f32,ne=[32768,512,1,1],k=1024,ties=0):                      375 runs -  3125.87 us/run -    67584 kB/run -   20.62 GB/s
  TOP_K(type=f32,ne=[65000,512,1,1],k=1024,ties=0):                      320 runs -  3890.73 us/run -   132048 kB/run -   32.37 GB/s
  TOP_K(type=f32,ne=[200000,512,1,1],k=1024,ties=0):                     168 runs -  6002.74 us/run -   402048 kB/run -   63.87 GB/s
  TOP_K(type=f32,ne=[4096,2048,1,1],k=1024,ties=0):                      205 runs -  5864.83 us/run -    40960 kB/run -    6.66 GB/s
  TOP_K(type=f32,ne=[8192,2048,1,1],k=1024,ties=0):                      228 runs -  7898.46 us/run -    73728 kB/run -    8.90 GB/s
  TOP_K(type=f32,ne=[16384,2048,1,1],k=1024,ties=0):                     122 runs - 10015.91 us/run -   139264 kB/run -   13.26 GB/s
  TOP_K(type=f32,ne=[32768,2048,1,1],k=1024,ties=0):                      96 runs - 12429.59 us/run -   270336 kB/run -   20.74 GB/s
  TOP_K(type=f32,ne=[65000,2048,1,1],k=1024,ties=0):                      80 runs - 15580.25 us/run -   528192 kB/run -   32.33 GB/s
  TOP_K(type=f32,ne=[200000,2048,1,1],k=1024,ties=0):                     42 runs - 24064.55 us/run -  1608192 kB/run -   63.73 GB/s
  Backend CPU: OK
2/2 backends passed
OK

GPU (CUB argsort)

$ ./bin/test-backend-ops -b CUDA0 perf -o "TOP_K"
ggml_cuda_init: found 1 CUDA devices (Total VRAM: 97247 MiB):
  Device 0: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition, compute capability 12.0, VMM: yes, VRAM: 97247 MiB
Testing 2 devices

Backend 1/2: CUDA0
  Device description: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition
  Device memory: 97247 MB (96640 MB free)

  TOP_K(type=f32,ne=[4096,1,1,1],k=1024,ties=0):               40960 runs -    26.71 us/run -       20 kB/run -    0.71 GB/s
  TOP_K(type=f32,ne=[8192,1,1,1],k=1024,ties=0):               16384 runs -   103.36 us/run -       36 kB/run -    0.33 GB/s
  TOP_K(type=f32,ne=[16384,1,1,1],k=1024,ties=0):              16384 runs -   103.04 us/run -       68 kB/run -    0.63 GB/s
  TOP_K(type=f32,ne=[32768,1,1,1],k=1024,ties=0):              16384 runs -   107.00 us/run -      132 kB/run -    1.18 GB/s
  TOP_K(type=f32,ne=[65000,1,1,1],k=1024,ties=0):              16384 runs -   117.12 us/run -      257 kB/run -    2.10 GB/s
  TOP_K(type=f32,ne=[200000,1,1,1],k=1024,ties=0):             16384 runs -   125.28 us/run -      785 kB/run -    5.98 GB/s
  TOP_K(type=f32,ne=[4096,512,1,1],k=1024,ties=0):              9831 runs -   136.90 us/run -    10240 kB/run -   71.33 GB/s
  TOP_K(type=f32,ne=[8192,512,1,1],k=1024,ties=0):              5463 runs -   260.37 us/run -    18432 kB/run -   67.51 GB/s
  TOP_K(type=f32,ne=[16384,512,1,1],k=1024,ties=0):                     2892 runs -   470.72 us/run -    34816 kB/run -   70.54 GB/s
  TOP_K(type=f32,ne=[32768,512,1,1],k=1024,ties=0):                      994 runs -  1126.98 us/run -    67584 kB/run -   57.19 GB/s
  TOP_K(type=f32,ne=[65000,512,1,1],k=1024,ties=0):                      510 runs -  2638.50 us/run -   132048 kB/run -   47.73 GB/s
  TOP_K(type=f32,ne=[200000,512,1,1],k=1024,ties=0):                     168 runs - 10834.52 us/run -   402048 kB/run -   35.39 GB/s
  TOP_K(type=f32,ne=[4096,2048,1,1],k=1024,ties=0):                     2460 runs -   441.08 us/run -    40960 kB/run -   88.56 GB/s
  TOP_K(type=f32,ne=[8192,2048,1,1],k=1024,ties=0):                      912 runs -  1124.20 us/run -    73728 kB/run -   62.54 GB/s
  TOP_K(type=f32,ne=[16384,2048,1,1],k=1024,ties=0):                     482 runs -  2145.99 us/run -   139264 kB/run -   61.89 GB/s
  TOP_K(type=f32,ne=[32768,2048,1,1],k=1024,ties=0):                     250 runs -  4459.60 us/run -   270336 kB/run -   57.81 GB/s
  TOP_K(type=f32,ne=[65000,2048,1,1],k=1024,ties=0):                     128 runs - 10552.70 us/run -   528192 kB/run -   47.73 GB/s
  TOP_K(type=f32,ne=[200000,2048,1,1],k=1024,ties=0):                     42 runs - 40201.88 us/run -  1608192 kB/run -   38.15 GB/s
  Backend CUDA0: OK
Backend 2/2: CPU
  Skipping
2/2 backends passed
OK

So it's not like GPU is always faster. My guess is that performance improvements will be mostly visible on platforms with slow CPU, slow CPU memory or slow CPU-GPU interconnect.

…checkpoint-all-opt Merge ROCm-specific optimizations: - PR ggml-org#26592: enable CUB path on HIP via hipCUB (top_k support) - ci : install hipCUB headers - perf: sweep quants for best performance - fix: q2_k mma - fix: q6_k mma - Placed MMA gateway over WMMA gateway - ggml : tune RDNA4 MMVQ warps for K-quants - server: persist checkpoints via sidecar file for hybrid/recurrent slot save/restore - heartbeat: restore CUDA noop kernel, add Vulkan back in, rename --no-sleep to --gpu-heartbeat

Geramy · 2026-08-05T19:35:50Z

@fairydreaming I would 100% agree with you on this. cub is generally more expensive, now we could create a top-k function if llama.cpp wants but right now its the same function as argsort, top-k specifies it doesn't need a sort function, that might help, might. but this PR is more about fixing a limitation then performance I have the other PR I closed that one for this one, the other one was a extension and optimization, we can say anything more than X top-k gets moved to cub otherwise run with the new element per thread packing top-k function, Just a thought anyway it sounds like more work and more "gotchas" for maintenance.

matt23654 · 2026-08-05T23:22:27Z

After the recent changes, all is working correctly for me now! (i.e. compiling works fine by default and the pragma stuff above is not necessary).

@Geramy I agree with your views on the purpose of the PR, as the op (for >1024) is currently missing from ROCm, which in some cases (eg ROCm over RPC) causes a crash, so adding it to the ROCm backend is important. It may be possible to improve on-device-topk performance by using a dedicated primitive like rocprim::topk, but I completely agree that should be some later PR to tune the performance.

fairydreaming · 2026-08-06T10:16:22Z

@matt23654 Do I understand correctly that rocprim::topk (looks similar to CUB DeviceTopK) operates only on a single array? We need something that handles multi-array top-k like DeviceBatchedTopK recently introduced in CCCL.

matt23654 · 2026-08-06T15:04:58Z

@fairydreaming It is my understanding that rocprim::topk and rocprim::topk_pairs provide similar functionality to cub::DeviceTopK. My thinking was that at the moment on a modern CUDA system, the preferred llama.cpp pathway uses cub::DeviceTopK::MaxPairs, so to mirror this the ROCm pathway could use rocprim::topk_pairs. These should still have the advantage of not having to sort outputs and so be faster than the argsort pathway?

I believe your expert knowledge that for dsv4 a batched topk would be even better. ROCm has recently added rocprim::segmented_topk and rocprim::segmented_topk_pairs which provide similar functionality to cub::DeviceBatchedTopK. Looks like they expect inputs in contiguous memory though. Could these be used for the optimizations you have in mind?

ROCm/rocm-libraries#3646 -- At the time of this PR, only ROCm unstable had implemented a faster algorithm, so only very new ROCm can possibly be faster.
ROCm/rocm-libraries#8843 -- The segmented functions are very new, I think just TheRock nightlies currently.

Geramy · 2026-08-06T15:09:40Z

@matt23654 I believe rocm 7.14 will be landing soon as well. @superm1 has a PR to bump to 7.14

Geramy · 2026-08-06T17:41:16Z

@fairydreaming How soon do you think we can merge this? I would like to get started on the next PR that jumps at me haha.

fairydreaming · 2026-08-06T18:09:42Z

@matt23654

@fairydreaming It is my understanding that rocprim::topk and rocprim::topk_pairs provide similar functionality to cub::DeviceTopK. My thinking was that at the moment on a modern CUDA system, the preferred llama.cpp pathway uses cub::DeviceTopK::MaxPairs, so to mirror this the ROCm pathway could use rocprim::topk_pairs. These should still have the advantage of not having to sort outputs and so be faster than the argsort pathway?

Yes but this solution is slow during prompt processing due to calling it iteratively for each row in a loop, hopefully replacing it with DeviceSegmentedTopK will fix that.

I believe your expert knowledge that for dsv4 a batched topk would be even better. ROCm has recently added rocprim::segmented_topk and rocprim::segmented_topk_pairs which provide similar functionality to cub::DeviceBatchedTopK. Looks like they expect inputs in contiguous memory though. Could these be used for the optimizations you have in mind?

Lightning indexer scores are contiguous, so it seems to be exactly what we need.

fairydreaming · 2026-08-06T18:39:06Z

@fairydreaming How soon do you think we can merge this? I would like to get started on the next PR that jumps at me haha.

@Geramy Hard to say, I'm not in charge here. People review things when they have free time, which is kind of random. But if you are bored already how about getting lightning indexer WMMA kernel working on ROCm? Should get AMD users a nice performance bump. You said something earlier about rocWMMA being deprecated, but isn't that only about rocWMMA repo being deprecated (moved to rocm-libraries) and not the API?

Geramy · 2026-08-06T18:45:11Z

But if you are bored already how about getting lightning indexer WMMA kernel working on ROCm? Should get AMD users a nice performance bump. You said something earlier about rocWMMA being deprecated

I was under the impression from some commits and comments that rocWMMA was removed from llama.cpp and replaced with MMA I’ll have to go back and double check that.

Yes I totally could work on lightning indexer, maybe if the above is true I just bring it to MMA instead.

fairydreaming · 2026-08-06T18:52:50Z

But if you are bored already how about getting lightning indexer WMMA kernel working on ROCm? Should get AMD users a nice performance bump. You said something earlier about rocWMMA being deprecated

I was under the impression from some commits and comments that rocWMMA was removed from llama.cpp and replaced with MMA I’ll have to go back and double check that.

Yes I totally could work on lightning indexer, maybe if the above is true I just bring it to MMA instead.

@Geramy Flash attention WMMA kernel was removed recently because there is MMA one that is faster. But it's not the case with lightning indexer kernel - still uses WMMA (but only for CUDA at the moment, ROCm falls back to slower vector kernel).

kyuz0 · 2026-08-06T18:52:53Z

I was under the impression from some commits and comments that rocWMMA was removed from llama.cpp and replaced with MMA I’ll have to go back and double check that.

I believe that is the case, I recall @JohannesGaessler working on this 6-7 months ago, in fact I don't use rocmWMMA builds any longer and recall benchmarking them against the new kernels and the new kernels performing on par or better.

IMbackK · 2026-08-06T19:08:58Z

yes rocwamma support in fattn_ext is gone.

…ind to test ds4-fast top-k functionality.

CISC

@IMbackK Update your review?

CISC · 2026-08-07T07:04:13Z

 #if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && CUDART_VERSION >= 11070
 #define USE_CUB
 #endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && CUDART_VERSION >= 11070

 #ifdef USE_CUB


Suggested change

#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && CUDART_VERSION >= 11070

#define USE_CUB

#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && CUDART_VERSION >= 11070

#ifdef USE_CUB

#ifdef GGML_CUDA_USE_CUB

CISC · 2026-08-07T07:04:23Z

 #ifdef USE_CUB
-#include <cub/cub.cuh>
 using namespace cub;
 #endif // USE_CUB


Suggested change

#endif // USE_CUB

#endif // GGML_CUDA_USE_CUB

..and update the other USE_CUB in this file...

github-actions Bot added ggml changes relating to the ggml tensor library for machine learning CUDA Related to the CUDA backend labels Aug 4, 2026

This was referenced Aug 4, 2026

CUDA: support ncols > 1024 in bitonic argsort/top-k The original bito… #26493

Closed

Eval bug: GGML_OP_TOP_K falls back to CPU on HIP/ROCm above ~3–4K context — 6.4× token-generation loss on DeepSeek-V4-Flash #26399

Open

fairydreaming reviewed Aug 4, 2026

View reviewed changes

Comment thread ggml/src/ggml-cuda/cub-compat.cuh Outdated

added missing newline.

aa58789

Geramy marked this pull request as ready for review August 4, 2026 19:18

Geramy requested review from a team and IMbackK as code owners August 4, 2026 19:18

Geramy requested a review from fairydreaming August 4, 2026 19:30

ci : install hipCUB headers

624eb53

fairydreaming requested a review from a team as a code owner August 5, 2026 06:38

ggerganov reviewed Aug 5, 2026

View reviewed changes

Comment thread ggml/src/ggml-cuda/vendors/hip.h Outdated

github-actions Bot added the devops improvements to build systems and github actions label Aug 5, 2026

cuda : add missing error checks

85d25f6

Update ggml/src/ggml-cuda/vendors/hip.h

5e5cd73

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>

IMbackK self-assigned this Aug 5, 2026

IMbackK requested changes Aug 5, 2026

View reviewed changes

Comment thread ggml/src/ggml-cuda/cub-compat.cuh Outdated

fairydreaming mentioned this pull request Aug 5, 2026

Eval bug: GLM-5.2 terrible performance on ROCm/HIP after Indexer PR #25407: prefill ~6x slower, load time ~40x longer #26445

Open

Made adjustments as requested, removing cub-compat and using vendor s…

f58f024

…pecific files for cub, dropped redundent includes and added a hip version check for hip cub inclusion, moved around a few includes, and added a test @fairydreaming originally had me run.

github-actions Bot added the testing Everything test related label Aug 5, 2026

Geramy requested a review from IMbackK August 5, 2026 19:36

fairydreaming reviewed Aug 6, 2026

View reviewed changes

Comment thread tests/test-backend-ops.cpp Outdated

disabled tests as requested by @fairydreaming but i want to leave beh…

d072d32

…ind to test ds4-fast top-k functionality.

Geramy requested a review from fairydreaming August 7, 2026 05:01

CISC approved these changes Aug 7, 2026

View reviewed changes

Uh oh!

Conversation

Geramy commented Aug 4, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Overview

Additional information

DeepSeek-V4-Flash UD-IQ1_S — 284B MoE, 78 GB

Qwen3.6-35B-A3B UD-Q4_K_XL

Requirements

Uh oh!

Uh oh!

fairydreaming commented Aug 4, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Neresco commented Aug 4, 2026

Uh oh!

Geramy commented Aug 4, 2026

Uh oh!

fairydreaming commented Aug 4, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Geramy commented Aug 4, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Neresco commented Aug 4, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

fairydreaming commented Aug 4, 2026

Uh oh!

Geramy commented Aug 4, 2026

Uh oh!

fairydreaming commented Aug 4, 2026

Uh oh!

remeh commented Aug 4, 2026

Uh oh!

Uh oh!

fairydreaming commented Aug 5, 2026

Uh oh!

Geramy commented Aug 5, 2026

Uh oh!

fairydreaming commented Aug 5, 2026

Uh oh!

kyuz0 commented Aug 5, 2026

Uh oh!

matt23654 commented Aug 5, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Geramy commented Aug 5, 2026

Uh oh!

kyuz0 commented Aug 5, 2026

Uh oh!

Uh oh!

Geramy commented Aug 5, 2026

Uh oh!

Geramy commented Aug 5, 2026

Uh oh!

fairydreaming commented Aug 5, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

CPU

GPU (CUB argsort)

Uh oh!

Geramy commented Aug 5, 2026

Uh oh!

matt23654 commented Aug 5, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

fairydreaming commented Aug 6, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

matt23654 commented Aug 6, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Geramy commented Aug 6, 2026

Uh oh!

Geramy commented Aug 6, 2026

Uh oh!

fairydreaming commented Aug 6, 2026

Uh oh!

fairydreaming commented Aug 6, 2026

Uh oh!

Geramy commented Aug 4, 2026 •

edited

Loading

fairydreaming commented Aug 4, 2026 •

edited

Loading

fairydreaming commented Aug 4, 2026 •

edited

Loading

Geramy commented Aug 4, 2026 •

edited

Loading

Neresco commented Aug 4, 2026 •

edited

Loading

matt23654 commented Aug 5, 2026 •

edited

Loading

fairydreaming commented Aug 5, 2026 •

edited

Loading

matt23654 commented Aug 5, 2026 •

edited

Loading

fairydreaming commented Aug 6, 2026 •

edited

Loading

matt23654 commented Aug 6, 2026 •

edited

Loading